
#include "minimo_gpu.h"
#include <assert.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <helper_cuda.h>
#include <device_functions.h> //_synchronizedThreads()
#include <stdlib.h>
#include <stdio.h>
#include <time.h>
#include <helper_image.h>
#include <sys/time.h>
#define SHARED_MEMX 64
#define SHARED_MEMY 64
extern __shared__ short unsigned int sumX[];

/**
 * add the kx rows and store the result in sumX
 * @param img
 * @param imgX image size in X
 * @param kx window size in X
 * @param ky window size in Y
 * @param dimX
 * @param dimY
 * @param pixelsPerThread integer that represents how match pixels a thread has to work with
 */
inline __device__ void sumRowMinimo(uchar4* img, int imgX, int kx,int ky,int dimX,int dimY, int pixelsPerThread){
	int kxMed = (int)kx/2;
	int blockY = (blockIdx.y*dimY);
	int blockX = (blockIdx.x*dimX);
	int length = dimX*dimY;
	uchar4 p;

	if(blockIdx.y!=0)
			blockY = (blockIdx.y*dimY)-blockIdx.y*ky;
	if(blockIdx.x!=0)
			blockX = (blockIdx.x*dimX)-blockIdx.x*kx;
	int idxX =dimX/pixelsPerThread;
	for(int X=threadIdx.x;X<dimX;X=X+idxX){
		int sum1=255,sum2=255,sum3=255;
		if(X>=kxMed && blockX+X<imgX-kxMed){
			for(int i=-kxMed;i<=kxMed;i++){
				p = img[(blockY+threadIdx.y)*imgX+blockX+X+i];
				sum1 = min(p.x,sum1);
				sum2 = min(p.y,sum2);
				sum3 = min(p.z,sum3);
			}
			sumX[threadIdx.y*dimX+X] = sum1;
			sumX[length+threadIdx.y*dimX+X] = sum2;
			sumX[2*length+threadIdx.y*dimX+X] = sum3;
		}
	}
}
/**
 * add ky columns and store the result in buf
 * @param buf image
 * @param imgX image size in X
 * @param imgY image size in Y
 * @param ky window size in Y
 * @param kx window size in X
 * @param dimX
 * @param dimY
 * @param pixelsPerThread integer that represents how match pixels a thread has to work with
 */
inline __device__ void sumColumnMinimo(uchar4* buf, int imgX, int imgY, int ky,int kx,int dimX,int dimY, int pixelsPerThread){

	int kyMed = (int)ky/2;
	int kxMed = (int)kx/2;
	int blockY = (blockIdx.y*dimY);
	int blockX = (blockIdx.x*dimX);
	int length = dimX*dimY;

	if(blockIdx.y!=0)
		blockY = (blockIdx.y*dimY)-blockIdx.y*ky;
	if(blockIdx.x!=0)
			blockX = (blockIdx.x*dimX)-blockIdx.x*ky;

	int idxX =dimX/pixelsPerThread;
	for(int X=threadIdx.x;X<dimX;X=X+idxX){
		int sum1=255,sum2=255,sum3=255;
		if(X>=kxMed && blockX+X<imgX-kxMed && threadIdx.y>=kyMed && threadIdx.y<dimY-kyMed && blockY+threadIdx.y<imgY-kyMed){
			uchar4 p;
			for(int j=-kyMed;j<=kyMed;j++){
				int idx = (threadIdx.y+j)*dimX+X;
				sum1 = min(sum1,sumX[idx]);
				sum2 = min(sum2,sumX[length+idx]);
				sum3 = min(sum3,sumX[2*length+idx]);
			}
			p.x=sum1;
			p.y=sum2;
			p.z=sum3;
			buf[(blockY+threadIdx.y)*imgX+blockX+X] = p;
		}
	}
}
/**
 * media filter in cuda
 * @param buf
 * @param img
 * @param imgX image size in X
 * @param imgY image size in Y
 * @param kx window size in X
 * @param ky window size in Y
 * @param dimX image size in X in the block
 * @param dimY image size in Y in the block
 * @param pixelsPerThread integer that represents how match pixels a thread has to work with
 */
__global__ void minimo(uchar4* buf,uchar4* img, int imgX, int imgY, int kx,int ky, int dimX, int dimY, int pixelsPerThread){

	sumRowMinimo(img,imgX,kx,ky,dimX,dimY,pixelsPerThread);
	__syncthreads();
	sumColumnMinimo(buf, imgX, imgY, ky,kx,dimX,dimY, pixelsPerThread);
}
/**
 * media filter
 * @param h_img image
 * @param width image size in X
 * @param height image size in Y
 * @param dims
 * @param kx window size in X
 * @param ky window size in Y
 * @param pixelsPerThread integer that represents how match pixels a thread has to work with
 * @return media filter image h_img
 */
uchar4* minimo_gpu(uchar4* h_img,int width, int height,int* dims, int kx, int ky, int pixelsPerThread){

	assert(h_img!=NULL && width!=0 && height!=0);

	uchar4* result = (uchar4*) calloc(sizeof(uchar4),width*height);
	//DEVICE
	uchar4* d_img;
	uchar4* d_buf;
	//ALOCAR MEMORIA HOST

	checkCudaErrors(cudaMalloc((void**) &(d_img), width*height*sizeof(uchar4)));
	checkCudaErrors(cudaMalloc((void**) &(d_buf), width*height*sizeof(uchar4)));

	//COPY IMG TO DEVICE
	checkCudaErrors(cudaMemcpy(d_img, h_img, width*height*sizeof(uchar4), cudaMemcpyHostToDevice));
	checkCudaErrors(cudaMemcpy(d_buf, d_img, width*height*sizeof(uchar4), cudaMemcpyDeviceToDevice));

	int mem = 3*SHARED_MEMX*SHARED_MEMY*sizeof(short unsigned int);

	dim3 dimGrid(width/(SHARED_MEMY-kx-1)+1,height/(SHARED_MEMY-ky-1)+1,1);
	dim3 threads(SHARED_MEMX/pixelsPerThread,min(SHARED_MEMY,height),1);

	minimo<<<dimGrid,threads,mem>>>(d_buf,d_img,width,height, kx,ky, SHARED_MEMX,SHARED_MEMY,pixelsPerThread);
	checkCudaErrors(cudaMemcpy(result, d_buf, width*height*sizeof(uchar4), cudaMemcpyDeviceToHost));
	printf("\n");
	printf("\n");
	printf("\n");
	/*printf("gpu\n");
	for(int j=0;j<height;j++){
		for(int i=0;i<width;i++){
			printf("%d\t",result[width*j+i].x);
		}
		printf("\n");
	}*/
	checkCudaErrors(cudaFree(d_img));
	checkCudaErrors(cudaFree(d_buf));
	return result;
}
